#include "kernel_operator.h"

using namespace AscendC;

constexpr int32_t BUFFER_NUM = 2; 

template<typename typeT>
class KernelReplicationPad2d {
public:
    __aicore__ inline KernelReplicationPad2d() {}
    __aicore__ inline void Init(GM_ADDR x, GM_ADDR y, int32_t lastdimlength, int32_t lastB2dimlength, int32_t totallength, int32_t tileDataNum, 
                                int32_t leftpad, int32_t rightpad, int32_t toppad, int32_t bottompad,int32_t outputlength, int32_t highdimcount)
    {
        this->lastdimlength = lastdimlength;
        this->lastB2dimlength = lastB2dimlength;
        this->totallength = totallength;
        this->tileDataNum = tileDataNum;
        this->leftpad = leftpad;
        this->rightpad = rightpad;
        this->toppad = toppad;
        this->bottompad = bottompad;
        this->outputlength = outputlength;
        this->highdimcount = highdimcount; 
        this->paddedlength = (lastdimlength + leftpad + rightpad) * (lastB2dimlength + toppad + bottompad);
        this->paddedline = lastdimlength + leftpad + rightpad;
        this->alignnum = 32 / sizeof(typeT);
        this->lastdimlengthalign = (lastdimlength + this->alignnum - 1) / this->alignnum * this->alignnum;

        // uint32_t alignnum = 32 / sizeof(typeT);
        // this->tilebufferlength = ((lastdimnum % alignnum) == 0) ? lastdimnum : (lastdimnum + alignnum - 1) / alignnum * alignnum;
        // uint32_t bufferlength = totallength + alignnum;

        xGm.SetGlobalBuffer((__gm__ typeT*)x, totallength);
        yGm.SetGlobalBuffer((__gm__ typeT*)y, outputlength);
        
        pipe.InitBuffer(inQueueX, BUFFER_NUM, this->tileDataNum * sizeof(typeT));
        pipe.InitBuffer(outQueueY, BUFFER_NUM, this->tileDataNum * sizeof(typeT));
        pipe.InitBuffer(outQueueYL, BUFFER_NUM, this->tileDataNum * sizeof(typeT));
        pipe.InitBuffer(outQueueYR, BUFFER_NUM, this->tileDataNum * sizeof(typeT));

    }        
    __aicore__ inline void Process()
    {
        for(int32_t k = 0; k < this->highdimcount; k++) {
            for(int32_t m = 0; m < this->lastB2dimlength; m++) {
                CopyIn(k, m);
                Compute();
                CopyOutT3(k, m);
            }
        }
    }
    __aicore__ inline void InitAlign(GM_ADDR x, GM_ADDR y, int32_t lastdimlength, int32_t lastB2dimlength, int32_t totallength, int32_t tileDataNum, 
                                int32_t leftpad, int32_t rightpad, int32_t toppad, int32_t bottompad,int32_t outputlength, int32_t highdimcount, int32_t outputlinealign)
    {
        this->lastdimlength = lastdimlength;
        this->lastB2dimlength = lastB2dimlength;
        this->totallength = totallength;
        this->tileDataNum = tileDataNum;
        this->leftpad = leftpad;
        this->rightpad = rightpad;
        this->toppad = toppad;
        this->bottompad = bottompad;
        this->outputlength = outputlength;
        this->highdimcount = highdimcount; 
        this->outputlinealign = outputlinealign;
        this->paddedlength = (lastdimlength + leftpad + rightpad) * (lastB2dimlength + toppad + bottompad);
        this->paddedline = lastdimlength + leftpad + rightpad;

        xGm.SetGlobalBuffer((__gm__ typeT*)x, totallength);
        yGm.SetGlobalBuffer((__gm__ typeT*)y, outputlength);
        
        pipe.InitBuffer(inQueueX, BUFFER_NUM, this->tileDataNum * sizeof(typeT));
        pipe.InitBuffer(outQueueY, BUFFER_NUM, this->tileDataNum * sizeof(typeT));

    }     
    __aicore__ inline void ProcessAlign()
    {
        for(int32_t k = 0; k < this->highdimcount; k++) {
            for(int32_t m = 0; m < this->lastB2dimlength; m++) {
                CopyIn(k, m);
                ComputeAlign();
                CopyOutAlign(k, m);
            }
        }
    }
    __aicore__ inline void ProcessAlignConnect()
    {
        this->connectnum = this->tileDataNum / this->paddedline;
        //int32_t loopcount = this->lastB2dimlength / this->connectnum;
        for(int32_t k = 0; k < this->highdimcount; ++k) {
            for(int32_t m = 0; m < this->lastB2dimlength; m++) {
                if(m == 0 || m == this->lastB2dimlength - 1) {
                    CopyInAlignConnect(k, m, 1);
                    ComputeAlignConnect(1);
                    CopyOutAlignConnect(k, m, 1);
                } else if(m + this->connectnum - 1 < this->lastB2dimlength - 1) {
                    CopyInAlignConnect(k, m, this->connectnum);
                    ComputeAlignConnect(this->connectnum);
                    CopyOutAlignConnect(k, m, this->connectnum);
                    m += this->connectnum - 1;
                } else {
                    int32_t connecttmp = this->lastB2dimlength - 1 - m;
                    CopyInAlignConnect(k, m, connecttmp);
                    ComputeAlignConnect(connecttmp);
                    CopyOutAlignConnect(k, m, connecttmp);
                    m += connecttmp - 1;
                }
            }
        }
    }

private:
    __aicore__ inline void CopyInAlignConnect(int32_t progress, int32_t line, int32_t connect)
    {
        int32_t startpos = (progress * this->lastdimlength * this->lastB2dimlength) + (line * this->lastdimlength);
        int32_t length = this->lastdimlength * connect;
        LocalTensor<typeT> xLocal = inQueueX.AllocTensor<typeT>();
        DataCopy(xLocal, xGm[startpos], length);
        inQueueX.EnQue(xLocal);

    }
    __aicore__ inline void ComputeAlignConnect(int32_t connect) {
        LocalTensor<typeT> xLocal = inQueueX.DeQue<typeT>();
        LocalTensor<typeT> yLocal = outQueueY.AllocTensor<typeT>();
        //Muls(yLocal, yLocal, typeT(0), this->paddedline * connect);
        for(int32_t k = 0; k < connect; k++) {
            int32_t starttmppos = k * this->paddedline;
            typeT lefttmp = xLocal.GetValue(k * this->lastdimlength);
            typeT righttmp = xLocal.GetValue(k * this->lastdimlength + this->lastdimlength - 1);
            Duplicate(yLocal[starttmppos], lefttmp, this->leftpad);
            Duplicate(yLocal[starttmppos + this->leftpad + this->lastdimlength], righttmp, this->rightpad);
            //Adds(yLocal[starttmppos], yLocal[starttmppos], lefttmp, this->leftpad);
            //Adds(yLocal[starttmppos + this->leftpad + this->lastdimlength], yLocal[starttmppos + this->leftpad + this->lastdimlength], righttmp, this->rightpad);
            Muls(yLocal[starttmppos + this->leftpad], xLocal[k * this->lastdimlength], typeT(1), this->lastdimlength);
        }
        inQueueX.FreeTensor(xLocal);
        outQueueY.EnQue<typeT>(yLocal);
    }
    __aicore__ inline void CopyOutAlignConnect(int32_t progress, int32_t line, int32_t connect) {
        LocalTensor<typeT> yLocal = outQueueY.DeQue<typeT>();

        int32_t startpos = progress * this->paddedlength;

        if(line == 0) {
            for(int32_t t = 0; t <= this->toppad; t++) {                
                // DataCopyPad(yGm[startpos], yLocal, copyParams);
                DataCopy(yGm[startpos], yLocal, this->paddedline);
                startpos += this->paddedline;
            }
            // uint16_t blockCount = this->toppad;
            // uint16_t blockLen = this->paddedline * sizeof(typeT) / 32;
            // uint16_t dstStride = this->paddedline * sizeof(typeT) / 32;
            // DataCopyParams copyParams{blockCount, blockLen, 0, dstStride};            
            // DataCopy(yGm[startpos], yLocal, copyParams);
        } else if(line == this->lastB2dimlength - 1) {
            startpos += (this->toppad + this->lastB2dimlength - 1) * this->paddedline;
            for(int32_t t = 0; t <= this->bottompad; t++) {                
                // DataCopyPad(yGm[startpos], yLocal, copyParams);
                DataCopy(yGm[startpos], yLocal, this->paddedline);
                startpos += this->paddedline;
            }
        } else {
            startpos += (this->toppad + line) * this->paddedline;
            //DataCopyPad(yGm[startpos], yLocal, copyParams);
            DataCopy(yGm[startpos], yLocal, this->paddedline * connect);
        }

        outQueueY.FreeTensor(yLocal);
    }    
    __aicore__ inline void CopyIn(int32_t progress, int32_t line)
    {
        // LocalTensor<typeT> xLocal = inQueueX.AllocTensor<typeT>();
        // DataCopy(xLocal, xGm[progress * length], this->tilebufferlength);
        // inQueueX.EnQue(xLocal);

        int32_t startpos = (progress * this->lastdimlength * this->lastB2dimlength) + (line * this->lastdimlength);  
        LocalTensor<typeT> xLocal = inQueueX.AllocTensor<typeT>();
        uint32_t blockLen = this->lastdimlength * sizeof(typeT);
        DataCopyExtParams copyParams{1, blockLen, 0, 0, 0}; 
        DataCopyPadExtParams<typeT> padParams{false, 0, 0, 0};
        DataCopyPad(xLocal, xGm[startpos], copyParams, padParams);
        inQueueX.EnQue(xLocal);

    }
    __aicore__ inline void Compute() {
        LocalTensor<typeT> xLocal = inQueueX.DeQue<typeT>();
        LocalTensor<typeT> yLocal = outQueueY.AllocTensor<typeT>();
        LocalTensor<typeT> ylLocal = outQueueYL.AllocTensor<typeT>();
        LocalTensor<typeT> yrLocal = outQueueYR.AllocTensor<typeT>();
        typeT lefttmp = xLocal.GetValue(0);
        typeT righttmp = xLocal.GetValue(this->lastdimlength - 1);
        Duplicate(ylLocal, lefttmp, this->leftpad);
        Duplicate(yrLocal, righttmp, this->rightpad);
        //yLocal = xLocal;
        Muls(yLocal, xLocal, typeT(1), this->lastdimlength);
        //Adds(yLocal, xLocal, typeT(0), this->lastdimlength);
        //DataCopy(yLocal, xLocal, this->lastdimlengthalign);
        inQueueX.FreeTensor(xLocal);
        outQueueY.EnQue<typeT>(yLocal);
        outQueueYL.EnQue<typeT>(ylLocal);
        outQueueYR.EnQue<typeT>(yrLocal);
    }
   __aicore__ inline void CopyOutT3(int32_t progress, int32_t line) {
        LocalTensor<typeT> yLocal = outQueueY.DeQue<typeT>();
        LocalTensor<typeT> ylLocal = outQueueYL.DeQue<typeT>();
        LocalTensor<typeT> yrLocal = outQueueYR.DeQue<typeT>();
        int32_t startpos = progress * this->paddedlength;
        uint32_t blockLenM = this->lastdimlength * sizeof(typeT);
        uint32_t blockLenL = this->leftpad * sizeof(typeT);
        uint32_t blockLenR = this->rightpad * sizeof(typeT);
        DataCopyExtParams copyParamsM{1, blockLenM, 0, 0, 0};
        DataCopyExtParams copyParamsL{1, blockLenL, 0, 0, 0};
        DataCopyExtParams copyParamsR{1, blockLenR, 0, 0, 0};
        if(line == 0) {
            for(int32_t t = 0; t <= this->toppad; t++) {                
                DataCopyPad(yGm[startpos], ylLocal, copyParamsL);
                DataCopyPad(yGm[startpos + this->leftpad], yLocal, copyParamsM);
                DataCopyPad(yGm[startpos + this->leftpad + this->lastdimlength], yrLocal, copyParamsR);
                startpos += this->paddedline;
            }
        } else if(line == this->lastB2dimlength - 1) {
            startpos += (this->toppad + this->lastB2dimlength - 1) * this->paddedline;
            for(int32_t t = 0; t <= this->bottompad; t++) {                
                DataCopyPad(yGm[startpos], ylLocal, copyParamsL);
                DataCopyPad(yGm[startpos + this->leftpad], yLocal, copyParamsM);
                DataCopyPad(yGm[startpos + this->leftpad + this->lastdimlength], yrLocal, copyParamsR);
                startpos += this->paddedline;
            }
        } else {
            startpos += (this->toppad + line) * this->paddedline;
            DataCopyPad(yGm[startpos], ylLocal, copyParamsL);
            DataCopyPad(yGm[startpos + this->leftpad], yLocal, copyParamsM);
            DataCopyPad(yGm[startpos + this->leftpad + this->lastdimlength], yrLocal, copyParamsR);
        }

        outQueueY.FreeTensor(yLocal);
        outQueueYL.FreeTensor(ylLocal);
        outQueueYR.FreeTensor(yrLocal);    
    }   
    __aicore__ inline void ComputeAlign() {
        LocalTensor<typeT> xLocal = inQueueX.DeQue<typeT>();
        LocalTensor<typeT> yLocal = outQueueY.AllocTensor<typeT>();

        typeT lefttmp = xLocal.GetValue(0);
        typeT righttmp = xLocal.GetValue(this->lastdimlength - 1);
        Duplicate(yLocal, lefttmp, this->leftpad);
        Duplicate(yLocal[this->leftpad + this->lastdimlength], righttmp, this->rightpad);
        Muls(yLocal[this->leftpad], xLocal, typeT(1), this->lastdimlength);
        // if(this->lastdimlength * sizeof(typeT) % 256 == 0) {
        //     uint64_t mask = 256 / sizeof(typeT);
        //     uint8_t repeatTimes = this->lastdimlength * sizeof(typeT) / 256;
        //     Copy(yLocal[this->leftpad], xLocal, mask, repeatTimes, { 1, 1, 8, 8 });
        // } else {
        //     Adds(yLocal[this->leftpad], xLocal, typeT(0), this->lastdimlength);
        // }
        inQueueX.FreeTensor(xLocal);
        outQueueY.EnQue<typeT>(yLocal);
    } 
    __aicore__ inline void CopyOutAlign(int32_t progress, int32_t line) {
        LocalTensor<typeT> yLocal = outQueueY.DeQue<typeT>();

        int32_t startpos = progress * this->paddedlength;
        uint32_t blockLen = this->paddedline * sizeof(typeT);
        DataCopyExtParams copyParams{1, blockLen, 0, 0, 0};

        if(line == 0) {
            for(int32_t t = 0; t <= this->toppad; t++) {                
                DataCopyPad(yGm[startpos], yLocal, copyParams);
                startpos += this->paddedline;
            }
        } else if(line == this->lastB2dimlength - 1) {
            startpos += (this->toppad + this->lastB2dimlength - 1) * this->paddedline;
            for(int32_t t = 0; t <= this->bottompad; t++) {                
                DataCopyPad(yGm[startpos], yLocal, copyParams);
                startpos += this->paddedline;
            }
        } else {
            startpos += (this->toppad + line) * this->paddedline;
            //DataCopyPad(yGm[startpos], yLocal, copyParams);
            DataCopy(yGm[startpos], yLocal, this->paddedline);
        }

        outQueueY.FreeTensor(yLocal);
    }    
    __aicore__ inline void CopyOut(int32_t progress, int32_t line)
    {
        LocalTensor<typeT> xLocal = inQueueX.DeQue<typeT>();
        LocalTensor<typeT> ylLocal = outQueueYL.AllocTensor<typeT>();
        LocalTensor<typeT> yrLocal = outQueueYR.AllocTensor<typeT>();
        typeT lefttmp = xLocal.GetValue(0);
        typeT righttmp = xLocal.GetValue(this->lastdimlength - 1);
        int32_t startpos = progress * this->paddedlength;
        if(line == 0) {
            Duplicate(ylLocal, lefttmp, this->leftpad);
            outQueueYL.EnQue<typeT>(ylLocal);
            CopyOutYLPad(startpos, this->leftpad, this->toppad + 1);

            Duplicate(yrLocal, righttmp, this->rightpad);
            outQueueYR.EnQue<typeT>(yrLocal);
            CopyOutYRPad(startpos + this->lastdimlength + this->leftpad, this->rightpad, this->toppad + 1); 

            startpos += this->leftpad;    
            for(int32_t t = 0; t <= this->toppad; t++) {
                uint32_t blockLen = this->lastdimlength * sizeof(typeT);
                DataCopyExtParams copyParams{1, blockLen, 0, 0, 0};
                DataCopyPad(yGm[startpos], xLocal, copyParams);
                startpos += this->paddedline;
            }
        } else if(line == this->lastB2dimlength - 1) {
            startpos += (this->toppad + this->lastB2dimlength - 1) * this->paddedline;
            Duplicate(ylLocal, lefttmp, this->leftpad);
            outQueueYL.EnQue<typeT>(ylLocal);
            CopyOutYLPad(startpos, this->leftpad, this->bottompad + 1);

            Duplicate(yrLocal, righttmp, this->rightpad);
            outQueueYR.EnQue<typeT>(yrLocal);
            CopyOutYRPad(startpos + this->lastdimlength + this->leftpad, this->rightpad, this->bottompad + 1); 

            startpos += this->leftpad;    
            for(int32_t t = 0; t <= this->bottompad; t++) {
                uint32_t blockLen = this->lastdimlength * sizeof(typeT);
                DataCopyExtParams copyParams{1, blockLen, 0, 0, 0};
                DataCopyPad(yGm[startpos], xLocal, copyParams);
                startpos += this->paddedline;
            }
        } else {
            startpos += (this->toppad + line) * this->paddedline;
            Duplicate(ylLocal, lefttmp, this->leftpad);
            outQueueYL.EnQue<typeT>(ylLocal);
            CopyOutYLPad(startpos, this->leftpad, 1);

            Duplicate(yrLocal, righttmp, this->rightpad);
            outQueueYR.EnQue<typeT>(yrLocal);
            CopyOutYRPad(startpos + this->lastdimlength + this->leftpad, this->rightpad, 1); 

            startpos += this->leftpad;    
            uint32_t blockLen = this->lastdimlength * sizeof(typeT);
            DataCopyExtParams copyParams{1, blockLen, 0, 0, 0};
            DataCopyPad(yGm[startpos], xLocal, copyParams);
            startpos += this->paddedline;
            
        }

        inQueueX.FreeTensor(xLocal);
    }
    __aicore__ inline void CopyOutYLPad(int32_t startpos, int32_t length, int32_t copytimes) {
        LocalTensor<typeT> yLocal = outQueueYL.DeQue<typeT>();
        for(int32_t t = 0; t < copytimes; t++) {
            uint32_t blockLen = length * sizeof(typeT);
            DataCopyExtParams copyParams{1, blockLen, 0, 0, 0};
            DataCopyPad(yGm[startpos], yLocal, copyParams); 
            startpos += this->paddedline;       
        }
        outQueueYL.FreeTensor(yLocal);
    }
    __aicore__ inline void CopyOutYRPad(int32_t startpos, int32_t length, int32_t copytimes) {
        LocalTensor<typeT> yLocal = outQueueYR.DeQue<typeT>();
        for(int32_t t = 0; t < copytimes; t++) {
            uint32_t blockLen = length * sizeof(typeT);
            DataCopyExtParams copyParams{1, blockLen, 0, 0, 0};
            DataCopyPad(yGm[startpos], yLocal, copyParams); 
            startpos += this->paddedline;       
        }
        outQueueYR.FreeTensor(yLocal);
    }
    __aicore__ inline void ComputeT6() {
        LocalTensor<typeT> xLocal = inQueueX.DeQue<typeT>();
        LocalTensor<typeT> ylLocal = outQueueYL.AllocTensor<typeT>();
        LocalTensor<typeT> yrLocal = outQueueYR.AllocTensor<typeT>();
        typeT lefttmp = xLocal.GetValue(0);
        typeT righttmp = xLocal.GetValue(this->lastdimlength - 1);
        Duplicate(ylLocal, lefttmp, this->leftpad);
        Duplicate(yrLocal, righttmp, this->rightpad);
        //Adds(yLocal, xLocal, typeT(0), this->lastdimlength);
        //DataCopy(yLocal, xLocal, this->lastdimlengthalign);
        inQueueX.EnQue<typeT>(xLocal);
        //outQueueY.EnQue<typeT>(yLocal);
        outQueueYL.EnQue<typeT>(ylLocal);
        outQueueYR.EnQue<typeT>(yrLocal);
    }
   __aicore__ inline void CopyOutT6(int32_t progress, int32_t line) {
        LocalTensor<typeT> xLocal = inQueueX.DeQue<typeT>();
        LocalTensor<typeT> ylLocal = outQueueYL.DeQue<typeT>();
        LocalTensor<typeT> yrLocal = outQueueYR.DeQue<typeT>();
        int32_t startpos = progress * this->paddedlength;
        uint32_t blockLenM = this->lastdimlength * sizeof(typeT);
        uint32_t blockLenL = this->leftpad * sizeof(typeT);
        uint32_t blockLenR = this->rightpad * sizeof(typeT);
        DataCopyExtParams copyParamsM{1, blockLenM, 0, 0, 0};
        DataCopyExtParams copyParamsL{1, blockLenL, 0, 0, 0};
        DataCopyExtParams copyParamsR{1, blockLenR, 0, 0, 0};
        if(line == 0) {
            for(int32_t t = 0; t <= this->toppad; t++) {                
                DataCopyPad(yGm[startpos], ylLocal, copyParamsL);
                DataCopyPad(yGm[startpos + this->leftpad], xLocal, copyParamsM);
                DataCopyPad(yGm[startpos + this->leftpad + this->lastdimlength], yrLocal, copyParamsR);
                startpos += this->paddedline;
            }
        } else if(line == this->lastB2dimlength - 1) {
            startpos += (this->toppad + this->lastB2dimlength - 1) * this->paddedline;
            for(int32_t t = 0; t <= this->bottompad; t++) {                
                DataCopyPad(yGm[startpos], ylLocal, copyParamsL);
                DataCopyPad(yGm[startpos + this->leftpad], xLocal, copyParamsM);
                DataCopyPad(yGm[startpos + this->leftpad + this->lastdimlength], yrLocal, copyParamsR);
                startpos += this->paddedline;
            }
        } else {
            startpos += (this->toppad + line) * this->paddedline;
            DataCopyPad(yGm[startpos], ylLocal, copyParamsL);
            DataCopyPad(yGm[startpos + this->leftpad], xLocal, copyParamsM);
            DataCopyPad(yGm[startpos + this->leftpad + this->lastdimlength], yrLocal, copyParamsR);
        }

        inQueueX.FreeTensor(xLocal);
        outQueueYL.FreeTensor(ylLocal);
        outQueueYR.FreeTensor(yrLocal);    
    }      
private:
    TPipe pipe;
    //TQueBind<TPosition::VECIN, TPosition::VECOUT, BUFFER_NUM> inQueueX;
    TQue<QuePosition::VECIN, BUFFER_NUM> inQueueX;
    TQue<QuePosition::VECOUT, BUFFER_NUM> outQueueY, outQueueYL, outQueueYR;
    GlobalTensor<typeT> xGm;
    GlobalTensor<typeT> yGm;

    int32_t lastdimlength;
    int32_t lastB2dimlength;
    int32_t totallength;
    int32_t tileDataNum;
    int32_t leftpad;
    int32_t rightpad;
    int32_t toppad; 
    int32_t bottompad;
    int32_t outputlength;
    int32_t highdimcount;
    int32_t paddedlength;
    int32_t paddedline;
    int32_t outputlinealign;
    int32_t connectnum;
    int32_t lastdimlengthalign;
    int32_t paddedlinealign;
    int32_t alignnum;
};

extern "C" __global__ __aicore__ void replication_pad2d(GM_ADDR x, GM_ADDR paddings, GM_ADDR y, GM_ADDR workspace, GM_ADDR tiling) {
    GET_TILING_DATA(t_data, tiling);
    KernelReplicationPad2d<DTYPE_X> op;
    if (TILING_KEY_IS(100)) {
        op.Init(x, y, t_data.lastdimlength, t_data.lastB2dimlength, t_data.totallength, t_data.tileDataNum,
                t_data.leftpad, t_data.rightpad, t_data.toppad, t_data.bottompad, t_data.outputlength, t_data.highdimcount);
        op.Process();
    } else if (TILING_KEY_IS(105)) {
        op.InitAlign(x, y, t_data.lastdimlength, t_data.lastB2dimlength, t_data.totallength, t_data.tileDataNum,
                t_data.leftpad, t_data.rightpad, t_data.toppad, t_data.bottompad, t_data.outputlength, t_data.highdimcount, t_data.outputlinealign);
        op.ProcessAlign();    
    } else if (TILING_KEY_IS(106)) {
        op.InitAlign(x, y, t_data.lastdimlength, t_data.lastB2dimlength, t_data.totallength, t_data.tileDataNum,
                t_data.leftpad, t_data.rightpad, t_data.toppad, t_data.bottompad, t_data.outputlength, t_data.highdimcount, t_data.outputlinealign);
        op.ProcessAlignConnect();         
    } else if (TILING_KEY_IS(200)){
        op.Init(x, y, t_data.lastdimlength, t_data.lastB2dimlength, t_data.totallength, t_data.tileDataNum,
                t_data.leftpad, t_data.rightpad, t_data.toppad, t_data.bottompad, t_data.outputlength, t_data.highdimcount);
        op.Process();
    }
}